	//.struct HeapChunk { .u64 Next; .u16 Magic; };
	//.struct Heap { .u64 FreeChunkChain; .u32 TotalChunks; };
	
	.global .u64 s_heap;
	
.func memoryHeapInitialize(.reg .u64 %heap, .reg .u32 %totalChunks)
{
	.reg .u64 %chunk;
	.reg .pred %p;
	
	// HeapChunk* chunk = (HeapChunk*)((__int8*)heap + sizeof(Heap));
	add.u64 %chunk, %heap, 16U; // sizeof(Heap) 
	
	// s_heap = (Heap*)heap;
	st.global.u64 [s_heap], %heap;

	// s_heap->FreeChunkChain = chunk;
	st.global.u64 [%heap + 0], %chunk;

	// s_heap->TotalChunks = totalChunks;
	st.global.u32 [%heap + 8], %totalChunks;

	// while (totalChunks-- > 1)
	setp.le.u32 %p, %totalChunks, 1U;
@%p bra $whileEnd;
$whileBegin:
	// chunk = chunk->Next = (HeapChunk*)((__int8*)chunk + sizeof(HeapChunk) + HEAPCHUNK_SIZE);
	add.u64 %chunk, %chunk, 1040U;
	st.global.u64 [%chunk + 0], %chunk;

	// chunk->Magic = HEAPCHUNK_MAGIC;
	st.global.u16 [%chunk + 8], 0x3412;
		
	// while (totalChunks-- > 1)
	sub.u32 %totalChunks, %totalChunks, 1U;
	setp.gt.u32 %p, %totalChunks, 1U;
@%p bra $whileBegin;
$whileEnd:
	
	// chunk->Next = nullT(HeapChunk*);
	st.global.u64 [%chunk + 0], 0U;
	
	// chunk->Magic = HEAPCHUNK_MAGIC;
	st.global.u16 [%chunk + 8], 0x3412;
	ret;
}

.func (.reg .u64 %obj) HeapGet()
{
	.reg .u64 %heap;
	.reg .u64 %chunk;
	.reg .u64 %r<1>;
	.reg .pred %p;
	
	// HeapChunk* chunk = s_heap->FreeChunkChain;
	ld.global.u64 %heap, [s_heap];
	ld.global.u64 %chunk, [%heap + 0];

	// if (chunk == null)
	//     throw(OutOfMemoryException);
	setp.eq.u64 %p, %chunk, 0U;
@%p	trap;

	// s_heap->FreeChunkChain = chunk->Next;
	ld.global.u64 %r0, [%chunk + 0];
	st.global.u64 [%heap + 0], %r0;

	// chunk->Next = nullT(HeapChunk*);
	st.global.u64 [%chunk + 0], 0U;

	// return (void*)((__int8*)chunk + sizeof(HeapChunk));
	sub.u64 %r0, %obj, 16;
	mov.u64 %obj, %chunk;
	ret;
}

.func memoryHeapFree(.reg .u64 %obj)
{
	.reg .u64 %heap;
	.reg .u64 %chunk;
	.reg .u64 %r<1>;
	.reg .pred %p;

	// HeapChunk* chunk = (HeapChunk*)((__int8*)obj - sizeof(HeapChunk));
	ld.global.u64 %heap, [s_heap];

	// if (chunk->Magic != HEAPCHUNK_MAGIC)
	//     throw(InvalidMagicException);
	//setp.eq.u64 %p, %chunkMagic, 0x3412;
@%p	trap;


	// chunk->Next = s_heap->FreeChunkChain;
	ld.global.u64 %r0, [%heap + 0];
	st.global.u64 [%obj + -16], %r0;

	sub.u64 %r0, %obj, 16;
	st.global.u64 [%heap + 0], %r0;
	ret;
}

/*
	.entry HeapGet
	{
	.reg .u64 %rd<6>;
	.loc	18	14	0
$LBB1_HeapGet:
	.loc	3	42	0
	ld.global.u64 	%rd1, [s_heap];
	ld.global.u64 	%rd2, [%rd1+8];
	.loc	3	46	0
	ld.global.u64 	%rd3, [%rd2+0];
	st.global.u64 	[%rd1+8], %rd3;
	.loc	3	47	0
	mov.u64 	%rd4, 0;
	st.global.u64 	[%rd2+0], %rd4;
	.loc	18	12	0
	exit;
$LDWend_HeapGet:
	} // HeapGet

	.entry HeapFree (
		.param .u64 __cudaparm_HeapFree_obj)
	{
	.reg .u64 %rd<6>;
	.loc	18	13	0
$LBB1_HeapFree:
	.loc	3	58	0
	ld.param.u64 	%rd1, [__cudaparm_HeapFree_obj];
	ld.global.u64 	%rd2, [s_heap];
	ld.global.u64 	%rd3, [%rd2+8];
	st.global.u64 	[%rd1+-16], %rd3;
	.loc	3	59	0
	sub.u64 	%rd4, %rd1, 16;
	st.global.u64 	[%rd2+8], %rd4;
	.loc	18	13	0
	exit;
$LDWend_HeapFree:
	} // HeapFree

*/